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摘 要 : 射电 脉冲 信号 在 传输 过 程 中 受到 星际 介质 的 影响 会 导致 轮廓 展 宽 和 变形 ， 在 研 
究 过 程 中 需要 对 信号 进行 消 色 散 处 理 。 本 文 设 计 并 实现 了 基于 零 找 贝 的 脉冲 星 数据 GPU 相干 
消 色 散 算法 ， 采 用 设备 内 存 映 射 以 消除 主机 到 设备 的 找 贝 开销 ， 利 用 CUDA 的 cuFFT 库 进 行 
多 BATCH 傅 里 叶 变 换 以 提高 DFT 效率 ,同时 采用 多 线程 实现 了 传递 函数 的 加 速 计算 。 实 验 结 
果 表 明 ， 与 传统 CPU 及 GPU 算法 相 比 ， 本 文 提出 的 算法 在 大 数据 量 时 表现 良好 。 
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脉冲 星 是 快速 旋转 的 中 子 星 ， 具 有 极 强 磁场 ， 于 1967 年 被 Bell 女士 发 现 。 自 此 ， 射 电 
脉冲 星 成 为 现代 天 文学 的 重要 研究 方向 之 一 。 以 射电 脉冲 星 作为 工具 , 可 以 开展 高 精度 计时 
和 授时 、 天 体 动力 学 和 天 体 测 量 、 强 场 下 的 引力 物理 、 太 阳 系 外 行星 、 星 系 和 星际 介质 、 超 
致密 物质 以 及 极端 环境 下 的 等 离子 物理 等 方面 的 研究 "。 近 年 来 , 利用 毫秒 脉冲 星 进行 引力 
波 探 测 、 脉 冲 星 导 航 等 项 目 逐 渐 兴 起 , 高 精度 脉冲 星 的 观测 技术 研究 成 为 射电 天 文学 的 热门 
问题 。 
脉冲 星 的 辐射 极其 微弱 ， 要 求 射 电 望远镜 有 具备 很 高 的 灵敏 度 。 建造 大 型 天 线 、 降 低 接收 
机 的 系统 噪声 和 增加 接收 机 的 频带 宽度 是 提高 灵敏 度 的 重要 方法 。 其 中 增 大 观测 带宽 是 经 济 、 
可 行 的 策略 ， 但 是 带宽 不 能 无 限 增 大 。 增 大 观测 带宽 ， 一 方面 会 给 后 续 处 理 带 来 巨大 压力 ， 
另 一 方面 星际 介质 会 对 脉冲 星 信 号 产生 严重 的 色散 影响 , 造成 脉冲 轮廓 的 展 宽 和 变形 , 严重 
时 甚至 可 能 导致 无 法 观测 到 脉冲 轮廓 ,因此 在 进行 脉冲 星相 关 研究 时 ， 必 须 对 脉冲 星 信号 进 
行 消 色散 处 理 

近年 来 ， 基 于 图 形 处 理 器 (Graphic Processing Unit, GPU) 的 并 行 计 算 技 术 已 成 为 高 性 
能 计算 领域 的 研究 热点 , 利用 图 形 处 理 器 可 大 大 提升 科学 分 析 、 仿真 等 方面 应 用 程序 的 运行 
WE, NVIDIA 推出 的 统一 计算 设备 架构 (Compute Unified Device Architecture, CUDA) 
可 使 GPU 解决 相对 复杂 的 问题 ”“， 为 计算 平台 的 构建 提供 了 保障 。 国 际 上 ， 使 用 CPU+GPU 的 
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异 构 计算 平台 进行 脉冲 信号 的 实时 处 理 已 成 为 主流 ,通过 使 用 GPU 对 观测 数据 进行 消 色散 处 
理 ， 可 降低 CPU 负载 ， 并 显著 提高 计算 系统 性 能 。 

本 文 针 对 脉冲 星 观 测 数据 ， 首 先 分 析 了 GPU 消 色 散 工作 原理 ， 然 后 通过 实验 研究 了 GPU 
消 色 散 算 法 的 性 能 瓶颈 ,最 后 设计 并 实现 了 基于 零 找 贝 的 GPU 相干 消 色散 算法 , 并 在 实验 中 
验证 了 算法 性 能 。 


1 色散 与 消 色 散 算法 

星际 介质 是 低温 的 等 离子 体 ， 脉 冲 星 的 电磁 辐射 通过 星际 介质 到 达 地 球 时 会 产生 色散 ， 
使 得 高 频 成 分 先 到 达 ， 低 频 成 分 后 到 达 。 色 散 是 脉冲 信和 号 的 最 重要 特征 ， 常 用 色散 度 
(Dispersion Measure，DM) 来 表示 ， 它 是 脉冲 星 与 地 球 的 视 向 距离 上 的 电子 密度 的 积分 ， 可 
由 公式 (1) 表 示 。 


DM = ff nedl (1) 


其 中 ，ne 是 电子 数 密度 ，d 为 脉冲 星 与 观测 望远镜 的 视 向 距离 。 
由 星际 介质 导致 频率 所 ,万 的 观测 时 延 At 如 公式 (2) 所 示 。 


At = kpm DM : (fr — fe? (2) 

Fh a a S78 2 A, MIETTA ERAF Be. EPA A a 
宽 为 BW 的 时 域 信号 由 滤波 器 分 成 若干 狭窄 通 道 ， 然 后 在 时 域 上 把 每 个 通道 信号 按 公 式 (2) 
提供 的 时 延 量 进行 平移 ， 最 后 将 所 有 通道 信号 幅度 累加 得 到 最 终 的 信号 序列 se 。 相 
于 消 色 散 将 星际 介质 对 脉冲 星 信号 的 影响 等 效 为 滤波 器 , 通常 的 做 法 是 将 望远镜 接收 的 信和 号 
进行 傅 里 叶 变 换 ， 然 后 乘 以 等 效 滤 波 器 传递 函数 的 反 函 数 即 chirp 函数 ,再 将 结果 进行 傅 里 
叶 道 变换 ， 即 可 得 到 时 域 上 的 脉冲 星 原 始 信和 号” 。 
星际 介质 传递 函数 可 分 解 为 幅度 响应 和 频率 响应 。 等 效 滤 波 器 需要 同时 对 信和 号 的 幅度 和 
相位 进行 修正 ， 由 此 可 得 离散 chirp 函数 的 表达 式 " (3) 。 
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chirp = T,H,* = 14) | exp| Cae ee) (8) 


GE for O<k<5 
fr = (k —N) 


N 
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了 为 taper 函数 ， 其 功能 是 对 脉冲 信号 的 幅度 进行 修正 。Hi ?是 脉冲 信号 传递 函数 相位 
响应 的 反 函 数 ， 其 功能 是 对 相位 进行 修正 。 其 中 N 为 离散 信号 的 频 域 点 数 ， 天 是 观测 的 中 心 
频率 ， 扩 是 频谱 中 第 k 点 的 频率 ，B 为 观测 带宽 。 
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图 1 异 构 程序 执行 模型 


Fig. 1 Heterogeneous program execution model 


N CPU+GPU 的 异 构 程序 执行 模型 如 图 1 所 示 , 其 核心 思想 是 将 串 行 代码 放置 在 主机 端 (CPU) 
© 上 执行 ， 而 将 并 行 代码 放置 在 设备 端 (GPU) 上 执行 ， 借 助 GPU 多 线程 来 减少 计算 开销 。 该 模 
N 型 假定 主机 和 设备 都 在 DRAM 中 保持 自己 独立 的 储存 空间 ， 分 别称 为 主机 存储 器 和 设备 存储 
CO 器 。 在 程序 执行 时 ， 需 要 将 预 处 理 数据 从 主机 存储 器 拷贝 到 设备 存储 器 ， 而 在 GPU 处 理 结束 
后 ， 还 要 将 运行 结果 拷贝 回 主机 存储 器 ， 这 造成 不 小 的 通信 开销 。 


float *H, *H MAP; 


: cudaDeviceProp prop; 
一 cudaGetDeviceProperties(&prop, 0) ; 
g if(!prop. canMapHostMemory) { 

exit (0); 
} 
cudaSetDeviceFlags (cudaDeviceMapHost) ; 
cudaHostAl loc (&H, nBytes, cudaHostAl locMapped) ; 
cudaHostGetDevicePointer (&H_MAP, H, 0) ; 
Kernel<<<gridSize, blockSize>>>(H_ MAP) ; 


图 2 零 找 贝 流程 


Fig.2 The process of Zero-copy 


FP LEJE GPU 直接 映射 主机 内 存 并 通过 PCIe 访问 而 无 需 进行 显 式 内 存 传输 ”。 它 需 
要 使 用 锁 页 映射 内 存 ， 核 函数 线程 可 直接 访问 主机 内 存 ， 从 而 减少 内 存 拷 贝 开 销 。 但 因为 数 
据 不 会 在 GPU 中 进行 缓存 , 所 以 锁 页 映射 内 存 只 能 被 读 取 和 写 入 一 次 , 多 次 读 写 会 降低 其 效 
率 ， 这 是 零 拷贝 的 局 限 性 。 

零 找 贝 的 代码 如 图 2 所 示 ， 其 中 cudaGetDeviceProperties () 返回 的 prop 通过 调用 
canMapHostMemory 来 检查 当前 GPU 是 否 文 持 将 主机 内 存 映射 到 设备 的 地 址 空间 。 通 过 使 用 
cudaDeviceMapHost 调用 cudaSetDeviceFlags 0 来 启动 锁 页 内 存 映 射 。 然 后 使 用 
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cudaHostAlloc 0 分 配 锁 页 映射 主机 内 存 ， 并 通过 函数 cudaHostGetDevicePointer () 获得 指 
向 映射 设备 地 址 空间 的 指针 。 在 零 找 贝 代码 中 , Kernel O 可 以 使 用 指针 HH MAP 引用 锁 页 主机 
内 存 ， 其 使 用 方式 与 调用 H_MAP 操作 设备 内 存 完全 相同 。 

3 基于 零 捞 贝 的 相干 消 色 散 算 法 


望远镜 在 接收 脉冲 星 信 号 时 会 受到 其 他 信号 的 干扰 ， 这 种 射频 干扰 (RFI) 可 定义 为 由 人 
类 活动 造成 的 不 恨 射 电信 号 。RFI 轻 则 会 降低 脉冲 星 数据 的 质量 ， 影 响 后 续 的 科学 研究 ， 重 
则 会 造成 脉 神 信号 丢失 ， 导 致 有 用 的 信号 完全 渡 没 在 噪声 中 。 因 此 必须 要 对 观测 数据 进行 
RFI 消减 。RFI 可 以 分 为 窄带 干扰 与 宽带 干扰 ， 全 球 定 位 系统 (GPS) 的 通信 信号 、 飞 机 通讯 信 
号 及 FM 广播 信号 等 都 属于 窜 带 干扰 ， 它 们 在 频 域 上 是 窄带 的 ， 但 因为 持续 时 间 长 ， 所 以 在 
时 域 上 长 期 存在 ; 雷电、 高 压 电缆、 电子 围栏 和 其 他 能 够 引起 瞬 态 电泳 的 信号 属于 宽带 干扰 ， 
它们 在 时 域 上 是 局 部 的 ， 但 会 污染 多 个 频 域 通 道 并 因此 表现 出 宽带 特性 。 
由 于 宽频 带 干 扰 在 频 域内 不 易 识 别 ， 而 在 时 域内 表现 为 短 时 强 脉冲 , 所 以 在 时 域 进行 宽 
iy RFI 的 消减 ， 具 体 方 法 是 : 
(1) 每 次 取 数 据 块 大 小 为 B， 由 n 个 点 组 成 ， 即 By, Bo -Bno 
(2) 求 数 据 块 值 的 均 方 根 (rms) 。 
(3) 使 用 均 方 根来 替代 B 中 值 超过 s .rms 的 数据 ， 其 中 s 是 设置 的 全 局 宽带 阔 值 。 


其 公式 描述 如 (4) 所 示 : 
n 2 
Bk = Ae, Bg 2 s: rms (4) 


mi TIAE EEEREN, 在 时 域 上 很 难 将 其 分 离 , 而 它 在 频 域 上 表现 为 菜 些 突 
出 的 尖峰 ， 所 以 在 频 域 进 行 窄带 RFI 的 消减 ， 具 体 方法 是 : 
(1) 将 经 过 宽带 RFI 消减 的 数据 块 B 变 换 至 频 域 并 用 F 表 示 ， 大 小 为 m, m = n/2+1. 
(2) WHFP RESA AE, Fo o En RELE, [Fol |x|， 然 后 求 模 值 的 平均 值 (mean)。 
(3) 生成 长 度 为 m 的 滤波 器 K， 如 果 | 玉 | 大 于 umean, $K = 0， WK, = 1， 其 中 v 是 设置 
的 全 局 窄带 阔 值 。 
(4) 将 K 与 F 相 乘 ， 得 到 消除 罕 带 RFI 后 的 频 域 序 列 。 


其 公式 描述 如 (5) 所 示 : 
F=F-K (5) 
其 中 
K = y |F;| >u-mean1<i<m 
i 1 else 
PilFi| 
mean = ==— 


m 

在 相干 消 色 散 算 法 中 , 需要 将 时 域 序列 变换 至 频 域 进行 处 理 , 这 就 要 求 对 数据 进行 离散 
傅 里 叶 变 换 (DFT) 。 为 了 使 傅 里 叶 变 换 更 加 高 效 ， 首 先 需要 对 原始 数据 进行 分 块 处 理 ， 然 后 
分 别 将 每 一 块 进行 DFT， 乘 以 星际 介质 的 传递 函数 进行 消 色 散 后 ， 再 进行 逆 变 换 (IDFT) ， 得 
到 最 终 的 时 间 序 列 。 

传统 的 CPU 使 用 FFTW 进行 DFT,FFTW 是 一 个 快速 计算 离散 健 里 叶 变 换 的 标准 C 程 序 集 ， 
是 目前 公认 速度 最 快 的 开源 FFT 变换 程序 ”。 在 NVIDIA GPU 上 进行 离散 傅 里 叶 变 换 需要 使 
用 cuFFT™, cuFFT 提供 了 经 过 大 量 优 化 和 广泛 测试 的 FFT 程序 库 ， 使 得 用 户 可 以 借助 GPU 


强大 的 计算 能 力 进行 快速 变换 。 


# define BATCH 10 
# define N 256 


ufftHandle plan; 


C 
C 
C 
C 


ufftComplex *host_in, *device_in; 
ufftReal *host_out, *device_out; 


ufftPlanld (&plan, N, CUFFT_R2C, BATCH) ; 


cufftExecR2C (plan, device in, device out); 


图 4 R2C DFT 核心 代码 


Fig.4 The core code of R2C DFT 


图 4 给 出 从 实数 到 复数 的 GPU DFT 变换 核心 代码 ,首先 定义 处 理 的 批 次 BATCH 和 每 一 批 


的 数据 量 N。 


然后 分 别 定义 cufftComplex 和 cufftReal 类 型 的 指针 ，host_ in 和 host_out 


分 别 代 表 主 机 端 (CPU) 输入 输出 数组 ， 而 device in 和 device out 代表 了 设备 端 (GPU) 的 输 


入 输出 数组 。 使 用 cui 
cufftPlanld 对 
少 次 一 维 变换 和 


每 次 变换 的 数据 量 ， 


fftHandle 定义 cuFFT 句柄 ， 即 一 个 执行 计划 。 接 着 使 用 一 
plan 进行 构造 ， 通 过 使 用 BATCH 和 N 进行 批 次 控制 来 告知 GPU 需要 进行 
实现 了 分 别 对 不 同 块 进行 DFT. 


EA 
多 


之 后 为 主机 端 数组 申请 锁 页 映射 内 存 ， 使 用 零 找 贝 完 成 内 存 到 显存 的 映射 ， 节 省 了 内 存 


拷贝 开销 ， 其 核心 代码 见 图 2。 内 存 分 配 完毕 
结果 写 入 device_out 中 。 值 得 注意 的 是 , cuFFT 在 执行 R2C 时 会 
， 那 么 得 到 的 复数 只 有 N/2+1 个 。 


即 如 果实 数 的 个 数 为 


后 ， 使 用 


cufftExecR2C 执行 傅 里 叶 变 换 并 将 
自动 去 除 DFT 的 对 称 部 分 ， 


为 了 最 大 化 指令 否 吐 量 , 核 函 数 中 的 数学 计算 全 部 使 用 CUDA 提供 的 
和 cosf 来 替换 传统 的 sin 和 cos， 前 者 能 够 在 
DFT 运行 过 程 中 采取 以 下 的 优化 策略 : 


外 ， 根 据 David Střelák 等 的 建议 汪 ， 在 


nly eR A EH sinf 
“影响 最 终结 果 的 前 提 下 加 快运 行 速度 。 此 


(1) 通过 对 数据 进行 截取 ， 保 证 数据 长 度 是 2 的 需 级 数 ， 加 快 了 运行 速度 。 


(2) 通过 使 用 单 精度 变换 ， 使 用 额外 空间 进行 DFT 结果 缓存 (out of place), ， 提 高 了 运 


行 效率 。 


TH 


(3) 通过 使 用 多 批 次 变换 进行 DFT， 而 不 是 重复 使 用 cufftPlanld， 缩 短 了 运行 时 间 。 


4 实验 结果 


算法 测试 所 使 用 的 开发 设备 如 表 1 tos, REFRA Windows 10， 开 发 工具 为 Visual 


Studio 2015、NVIDIA Visual Profiler, H 


F 发 语言 为 C、CUDA C, Python. 


表 1 开发 设备 


Tab. 1 Development equipment 


Name Device Type Storage Frequency 
CPUA Intel Core i7-4720HQ 8GB 2. 6GHZ 
GPUA NVIDIA GeForce GTX 960M 4GB 1. 1GHZ 
CPUB Intel Core i9-9900K 64GB 3. 6GHZ 
GPUB NVIDIA GeForce RTX 2080 SUPER 8GB 1. 65GHZ 


使 用 真实 的 脉冲 星 基带 数据 进行 算法 测试 


。 数 据 格式 为 双 极 化 8 比特 采样 的 PSRDADA, 
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观测 源 为 J0437-4715， 观 测 信息 如 表 2 所 示 。 基 带 数据 记录 时 间 约 为 8s， 大 小 为 12. 86B- 


表 2 J0437-4715 观测 信息 
Tab. 2 J0437-4715 Observation information 


J0437-4715 Information 
Period 0. 005757451936712637s 
DM 2. 64476 
Center frequency 1382MHZ 
Observation bandwidth 400MHZ 
Telescope Parkes 


为 验证 本 文 算法 的 正确 性 ， 进行 以 下 实验 : 首先 对 11. 72GB 的 原始 数据 i 
轮廓 图 作为 对 照 。 然 后 固定 数据 块 大 小 为 16MB， 使 用 本 文 算法 分 别处 理 50、 


THE, EM 
250, 750 个 数 


据 块 ， 其 数据 量 分 别 为 0.78GB、3. 91GB, 11. 72GB， 最 后 生成 不 同情 况 下 的 轮廓 图 。 实 验 结 
果 如 图 5 所 示 ， 各 子 图 横 轴 表示 脉冲 相位 ， 纵 轴 表 示 归 一 化 振幅 。(a) 代表 原始 数据 折 县 未 


进行 消 色散 ，(b) 、(c) (O 分 别 代表 数据 块 个 数 为 50、 数 据 块 个 数 为 250、 
750 时 消 色散 后 得 到 的 脉冲 轮廓 。 


Pulse phase - Normalized amplitude relationship 
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图 5 不 同情 况 下 的 脉冲 相位 - 归 一 化 振幅 关系 图 


数据 块 个 数 为 


Fig.5 Pulse phase -Normalized amplitude relationship in different situations 


对 比 图 5 中 的 (a)、(d) 可 得 ,本 文 算法 可 以 正确 的 对 脉冲 信号 进行 修正 ， 


脉冲 轮廓 。 通 过 与 新 疆 天 文 台数 据 中 心 的 J0437-4715 轮廓 进行 比 对 ， 证 明 所 得 轮廓 的 正确 
性 。 MEEK Cb). (c)s Cd) 可 知 ， 随 着 数据 块 数量 的 不 断 增 加 ， 所 得 脉冲 轮廓 信 品 比 逐 渐 提 高 。 
为 验证 CPU 算法 、 传统 GPU 算法 以 及 本 文 算 法 (ZGPU) 在 不 同 数据 量 下 的 执行 速度 , 进行 


完整 的 展现 出 


以 下 实验 : 首先 固定 数据 块 大 小 为 8M， 将 每 次 处 理 的 数据 块 数 量 分 别 设置 为 10、50、100、 
200、400、800、1500， 对 应 数据 量 为 0. 08GB、0. 39GB、0. 78GB、1. 56GB、3. 13GB、6. 25GB、 
11. 726GB， 然 后 令 三 种 方法 分 别 在 每 个 数据 量 下 运行 十 次 ， 最 后 取 平 均值 作为 运行 时 间 。 此 
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外 ， 为 验证 设备 性 能 对 算法 的 影响 ， 本 实验 采用 两 套 设备 进行 测试 ， 设 备 型 号 如 表 1 所 示 。 
实验 结果 如 表 3 所 示 ， 时 间 单 位 为 s。 
表 3 实验 结果 
Tab. 3 Experimental results 


Equipment Model Algorithm 10 50 100 200 400 800 1500 
CPU CPUA CP 9.213 42.351 84.024 165.608 330.420 659.481 1259.977 
CPUB GP 5.864 26.270 51.756 102.566 205.499 408.408 768.446 
GPU GPUA GP 2. 142 6. 351 11. 799 23. 301 47. 097 94. 287 180. 296 
ZGP 2.054 6.053 11. 230 21.734 43. 100 87. 178 166. 208 
GPUB GP 0.830 2.248 4. 018 7. 561 14. 635 28. 794 53. 603 
ZGP 0.808 2.130 3.781 7.077 13. 680 26. 859 49. 945 


CPUB-CPU vs GPUB-ZGPU 
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图 6 CPUB-CPU 与 GPUB-ZGPU 数据 量 -运行 时 间 关 系 图 
Fig. 6 CPUB-CPU vs GPUB-ZGPU Data size-Running time relationship 


GPUB-GPU vs GPUB-ZGPU 
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图 7 GPUB-GPU 与 GPUB-ZGPU 数据 量 - 运 行 时 间 关 系 图 
Fig. 7 GPUB-GPU vs GPUB-ZGPU Data size-Running time relationship 
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ZGPU vs GPUB-ZGPU 
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图 8 GPUA-ZGPU 与 GPUB-ZGPU 数据 量 - 运 行 时 间 关 系 图 


Fig. 8 GPUA-ZGPU vs GPUB-ZG 


PU Data size-Running time relationship 


本 文采 用 Model-Al gorithm 的 方式 对 算法 进行 讨论 , 例如 GPUA-GPU 即 代表 运行 在 GPUA 
上 的 传统 GPU 算法 。 从 表 3 中 抽取 CPUB-CPU 与 GPUB-ZGPU 做 折线 图 , 得 到 图 6; 取 GPUB-GPU 
与 GPUB-ZGPU 做 折线 图 ， 得 到 图 7; 取 GPUA-ZGPU 和 GPUB-ZGPU 得 到 图 8。 

结合 表 3 分 析 图 6 可 知 : 两 种 算法 的 执行 时 间 都 随 着 数据 量 增加 而 线性 增长 ,但 是 传统 


CPU 算法 的 平均 运行 时 间 远 大 于 本 文 算法 。 


在 数据 大 小 为 11. 726B 时 ， 传 统 CPU 算法 需要 


768. 446s 才能 完成 消 色 散 ， 而 本 文 算法 仅 需 49. 945s， 速 度 提 升 了 约 15 倍 。 由 此 可 得 ， 本 


文 算法 在 速度 上 优 于 传统 CPU 算法 。 


结合 表 3 分 析 图 7 可 知 : 两 种 算法 的 执行 时 间 都 随 着 数据 量 增 加 而 线性 增长 , 在 数据 量 
较 小 时 ， 本 文 算法 与 传统 GPU 算法 的 差距 不 明显 , 但 随 着 数据 量 的 增 大 ,两 种 算法 的 时 间 差 


也 在 线性 增长 。 当 数据 量 分 别 为 400、800、 


1500 时 ,两 种 算法 的 时 间 差 为 0.949s、1. 935s、 


3. 658s， 即 数据 量 增加 2 倍 ， 时 间 差 也 增加 2 倍 。 由 此 可 得 ， 本 文 算法 在 数据 量 较 大 时 优 于 
传统 GPU 算法 ， 数 据 量 越 大 ， 本 文 算法 的 优势 就 越 明显 。 


结合 表 3 分 析 图 8 可 知 : 本 文 算法 在 


新 架构 GPU 上 的 平均 运行 时 间 明 显 比 之 前 的 GPU 


短 。 在 数据 量 为 1500 WY, 本 文 算法 在 GPUA 上 需要 运行 166. 208s, , 而 在 GPUB 上 仅 需 49. 945s, 
速度 提高 了 约 3 倍 。 由 此 可 得 ， 硬 件 性 能 是 影响 本 文 算法 的 重要 因素 。 


5 结论 


本 文 设计 实现 了 基于 零 捞 贝 的 GPU 脉 六 


星相 干 消 色 散 算法 。 通 过 合理 构造 核 函数 实现 了 


最 大 化 并 行 执行 ， 采 取 零 拷贝 技术 提高 了 内 存 吞 吐 量 ， 使 用 CUDA 内 置 计算 函数 优化 了 指令 
调用 。 通 过 在 DFT 中 对 cuFFT 进行 优化 , 提高 了 其 运行 效率 。 结 合 真 实数 据 测试 表明 本 文 实 
现 的 算法 优 于 传统 的 CPU 算法 ， 且 在 数据 量 较 大 时 优 于 传统 的 GPU 算法 。 基 于 零 找 贝 的 GPU 


消 色 散 算法 为 高 效 实 现 脉冲 星 数据 处 理 相关 研究 提供 了 一 种 新 的 解决 思路 。 实 验 数据 及 源 代 


码 可 在 新 疆 天 文 台数 据 中 心 获取 。 


1 http://data.xao.ac.cn/static/zerocopy.rar 
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Abstract:The radio pulse signal is affected by the interstellar medium during the 
transmission process, which will cause the profile to be broadened and deformed. In the research 
process, the signal needs to be de-dispersed. This paper designed and implemented a 
zero-copy-based GPU coherent de-dispersion algorithm for pulsar data, used device memory 
mapping to eliminate the copy cost from host to device, and used CUDA's cuFFT library to 
perform multi-BATCH Fourier transform to improve DFT efficiency. Multithreading was used to 
realize the accelerated calculation of the transfer function. The experimental results show that, 
compared with traditional CPU and GPU algorithms, the algorithm proposed in this paper 
performs well in large amounts of data. 


Key words: zero-copy, coherent dedispersion, GPU, CUDA, cuFFT 


